

#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable
#define ACCUM_FLT4 float4
#define FLT float
#define FLT2 float2
#define FLT3 float3
#define FLT4 float4
#define TO_FLT4 convert_float4
#define TO_ACCUM_TYPE convert_float4
#define TO_ACCUM_FLT convert_float
#define READ_IMAGE read_imagef
#define WRITE_IMAGE write_imagef
__constant sampler_t
smp_edge = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
__constant sampler_t
smp_none = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;
__constant sampler_t
smp_zero = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
#define ACCUM_FLT float
constant ACCUM_FLT Bt[36] = {
        1.0000000000f, 0.0000000596f, -2.5000004768f, -0.0000001192f, 1.0000001192f, 0.0000000000f,
        0.0000000000f, 0.9428091049f, 1.3333333731f, -0.4714044929f, -0.6666667461f, 0.0000000000f,
        0.0000000000f, -0.9428090453f, 1.3333334923f, 0.4714045525f, -0.6666667461f, 0.0000000000f,
        0.0000000000f, -0.1178511307f, -0.0833333358f, 0.2357022613f, 0.1666666865f, 0.0000000000f,
        0.0000000000f, 0.1178511456f, -0.0833333507f, -0.2357022911f, 0.1666666865f, 0.0000000000f,
        0.0000000000f, 1.0000000000f, -0.0000000596f, -2.5000000000f, 0.0000000000f, 1.0000000000f,
};

__kernel void main_function(
        __read_only image2d_t bt_tex2d,
        __write_only image2d_t dst_tensor_image2d,
        __read_only image2d_t src_tensor_image2d,
        int4 shared_int4_0,
        int4 shared_int4_1)
{
    int DST_X = get_global_id(0);
    int DST_Y = get_global_id(1);
    int DST_Z = get_global_id(2);
    if (DST_X >= shared_int4_0.x || DST_Y >= 6 || DST_Z >= shared_int4_0.y)
    {
        return;
    }
    int tile_x = (DST_X % shared_int4_0.z) * 4;
    int tile_y = (DST_X / shared_int4_0.z) * 4;
    ACCUM_FLT4 I0, I1, I2, I3, I4, I5;
    ACCUM_FLT bt_ar[6];
    ACCUM_FLT4 t0 = TO_ACCUM_TYPE(read_imagef(bt_tex2d, smp_none, (int2)(DST_Y * 2 + 0, 0)));
    ACCUM_FLT4 t1 = TO_ACCUM_TYPE(read_imagef(bt_tex2d, smp_none, (int2)(DST_Y * 2 + 1, 0)));
    DST_Y *= 6;
    bt_ar[0] = t0.x;
    bt_ar[1] = t0.y;
    bt_ar[2] = t0.z;
    bt_ar[3] = t0.w;
    bt_ar[4] = t1.x;
    bt_ar[5] = t1.y;
    {
        int yc = tile_y + shared_int4_0.w;
        ACCUM_FLT bt = bt_ar[0];
        ACCUM_FLT4 src0 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 0), (yc) * shared_int4_1.y + (DST_Z)));
        I0 = bt * src0;
        ACCUM_FLT4 src1 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 1), (yc) * shared_int4_1.y + (DST_Z)));
        I1 = bt * src1;
        ACCUM_FLT4 src2 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 2), (yc) * shared_int4_1.y + (DST_Z)));
        I2 = bt * src2;
        ACCUM_FLT4 src3 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 3), (yc) * shared_int4_1.y + (DST_Z)));
        I3 = bt * src3;
        ACCUM_FLT4 src4 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 4), (yc) * shared_int4_1.y + (DST_Z)));
        I4 = bt * src4;
        ACCUM_FLT4 src5 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 5), (yc) * shared_int4_1.y + (DST_Z)));
        I5 = bt * src5;
    }
    {
        int yc = tile_y + shared_int4_0.w + (1);
        ACCUM_FLT bt = bt_ar[1];
        ACCUM_FLT4 src0 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 0), (yc) * shared_int4_1.y + (DST_Z)));
        I0 += bt * src0;
        ACCUM_FLT4 src1 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 1), (yc) * shared_int4_1.y + (DST_Z)));
        I1 += bt * src1;
        ACCUM_FLT4 src2 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 2), (yc) * shared_int4_1.y + (DST_Z)));
        I2 += bt * src2;
        ACCUM_FLT4 src3 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 3), (yc) * shared_int4_1.y + (DST_Z)));
        I3 += bt * src3;
        ACCUM_FLT4 src4 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 4), (yc) * shared_int4_1.y + (DST_Z)));
        I4 += bt * src4;
        ACCUM_FLT4 src5 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 5), (yc) * shared_int4_1.y + (DST_Z)));
        I5 += bt * src5;
    }
    {
        int yc = tile_y + shared_int4_0.w + (2);
        ACCUM_FLT bt = bt_ar[2];
        ACCUM_FLT4 src0 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 0), (yc) * shared_int4_1.y + (DST_Z)));
        I0 += bt * src0;
        ACCUM_FLT4 src1 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 1), (yc) * shared_int4_1.y + (DST_Z)));
        I1 += bt * src1;
        ACCUM_FLT4 src2 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 2), (yc) * shared_int4_1.y + (DST_Z)));
        I2 += bt * src2;
        ACCUM_FLT4 src3 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 3), (yc) * shared_int4_1.y + (DST_Z)));
        I3 += bt * src3;
        ACCUM_FLT4 src4 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 4), (yc) * shared_int4_1.y + (DST_Z)));
        I4 += bt * src4;
        ACCUM_FLT4 src5 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 5), (yc) * shared_int4_1.y + (DST_Z)));
        I5 += bt * src5;
    }
    {
        int yc = tile_y + shared_int4_0.w + (3);
        ACCUM_FLT bt = bt_ar[3];
        ACCUM_FLT4 src0 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 0), (yc) * shared_int4_1.y + (DST_Z)));
        I0 += bt * src0;
        ACCUM_FLT4 src1 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 1), (yc) * shared_int4_1.y + (DST_Z)));
        I1 += bt * src1;
        ACCUM_FLT4 src2 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 2), (yc) * shared_int4_1.y + (DST_Z)));
        I2 += bt * src2;
        ACCUM_FLT4 src3 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 3), (yc) * shared_int4_1.y + (DST_Z)));
        I3 += bt * src3;
        ACCUM_FLT4 src4 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 4), (yc) * shared_int4_1.y + (DST_Z)));
        I4 += bt * src4;
        ACCUM_FLT4 src5 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 5), (yc) * shared_int4_1.y + (DST_Z)));
        I5 += bt * src5;
    }
    {
        int yc = tile_y + shared_int4_0.w + (4);
        ACCUM_FLT bt = bt_ar[4];
        ACCUM_FLT4 src0 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 0), (yc) * shared_int4_1.y + (DST_Z)));
        I0 += bt * src0;
        ACCUM_FLT4 src1 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 1), (yc) * shared_int4_1.y + (DST_Z)));
        I1 += bt * src1;
        ACCUM_FLT4 src2 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 2), (yc) * shared_int4_1.y + (DST_Z)));
        I2 += bt * src2;
        ACCUM_FLT4 src3 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 3), (yc) * shared_int4_1.y + (DST_Z)));
        I3 += bt * src3;
        ACCUM_FLT4 src4 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 4), (yc) * shared_int4_1.y + (DST_Z)));
        I4 += bt * src4;
        ACCUM_FLT4 src5 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 5), (yc) * shared_int4_1.y + (DST_Z)));
        I5 += bt * src5;
    }
    {
        int yc = tile_y + shared_int4_0.w + (5);
        ACCUM_FLT bt = bt_ar[5];
        ACCUM_FLT4 src0 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 0), (yc) * shared_int4_1.y + (DST_Z)));
        I0 += bt * src0;
        ACCUM_FLT4 src1 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 1), (yc) * shared_int4_1.y + (DST_Z)));
        I1 += bt * src1;
        ACCUM_FLT4 src2 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 2), (yc) * shared_int4_1.y + (DST_Z)));
        I2 += bt * src2;
        ACCUM_FLT4 src3 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 3), (yc) * shared_int4_1.y + (DST_Z)));
        I3 += bt * src3;
        ACCUM_FLT4 src4 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 4), (yc) * shared_int4_1.y + (DST_Z)));
        I4 += bt * src4;
        ACCUM_FLT4 src5 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_x + shared_int4_1.x + 5), (yc) * shared_int4_1.y + (DST_Z)));
        I5 += bt * src5;
    }
    {
        FLT4 r0 = TO_FLT4(I0 + Bt[2] * I2 + Bt[4] * I4);
        write_imagef(dst_tensor_image2d, (int2)((DST_X), (DST_Y) * shared_int4_0.y + (DST_Z)), r0);;
        DST_Y++;
    }
    {
        FLT4 r0 = TO_FLT4(Bt[7] * I1 + Bt[8] * I2 + Bt[9] * I3 + Bt[10] * I4);
        write_imagef(dst_tensor_image2d, (int2)((DST_X), (DST_Y) * shared_int4_0.y + (DST_Z)), r0);;
        DST_Y++;
    }
    {
        FLT4 r0 = TO_FLT4(Bt[13] * I1 + Bt[14] * I2 + Bt[15] * I3 + Bt[16] * I4);
        write_imagef(dst_tensor_image2d, (int2)((DST_X), (DST_Y) * shared_int4_0.y + (DST_Z)), r0);;
        DST_Y++;
    }
    {
        FLT4 r0 = TO_FLT4(Bt[19] * I1 + Bt[20] * I2 + Bt[21] * I3 + Bt[22] * I4);
        write_imagef(dst_tensor_image2d, (int2)((DST_X), (DST_Y) * shared_int4_0.y + (DST_Z)), r0);;
        DST_Y++;
    }
    {
        FLT4 r0 = TO_FLT4(Bt[25] * I1 + Bt[26] * I2 + Bt[27] * I3 + Bt[28] * I4);
        write_imagef(dst_tensor_image2d, (int2)((DST_X), (DST_Y) * shared_int4_0.y + (DST_Z)), r0);;
        DST_Y++;
    }
    {
        FLT4 r0 = TO_FLT4(Bt[31] * I1 + Bt[33] * I3 + I5);
        write_imagef(dst_tensor_image2d, (int2)((DST_X), (DST_Y) * shared_int4_0.y + (DST_Z)), r0);;
        DST_Y++;
    }
}

#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable
#define ACCUM_FLT4 float4
#define FLT float
#define FLT2 float2
#define FLT3 float3
#define FLT4 float4
#define TO_FLT4 convert_float4
#define TO_ACCUM_TYPE convert_float4
#define TO_ACCUM_FLT convert_float
#define READ_IMAGE read_imagef
#define WRITE_IMAGE write_imagef
__constant sampler_t
smp_edge = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
__constant sampler_t
smp_none = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;
__constant sampler_t
smp_zero = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
#define ACCUM_FLT float
constant ACCUM_FLT At[24] = {
        1.0000000000f, 1.0000000000f, 1.0000000000f, 1.0000000000f, 1.0000000000f, 0.0000000000f,
        0.0000000000f, 0.7071067691f, -0.7071067691f, 1.4142135382f, -1.4142135382f, 0.0000000000f,
        0.0000000000f, 0.4999999702f, 0.4999999702f, 1.9999998808f, 1.9999998808f, 0.0000000000f,
        0.0000000000f, 0.3535533845f, -0.3535533845f, 2.8284270763f, -2.8284270763f, 1.0000000000f,
};

__kernel void main_function(
        __read_only image2d_t at_tex2d,
        __read_only image2d_t biases_tex2d,
        __write_only image2d_t dst_tensor_image2d,
        __read_only image2d_t src_tensor_image2d,
        int4 shared_int4_0,
        int4 shared_int4_1)
{
    int tile_id = get_global_id(0);
    int DST_Y = get_global_id(1);
    int DST_Z = get_global_id(2);
    int tile_x = (tile_id % shared_int4_0.x) * 4;
    int tile_y = (tile_id / shared_int4_0.x) * 4 + DST_Y;
    if (tile_x >= shared_int4_0.y || tile_y >= shared_int4_0.z || DST_Z >= shared_int4_0.w)
    {
        return;
    }
    ACCUM_FLT4 I0, I1, I2, I3, I4, I5;
    ACCUM_FLT at_ar[6];
    ACCUM_FLT4 t00 = TO_ACCUM_TYPE(read_imagef(at_tex2d, smp_none, (int2)(DST_Y * 2 + 0, 0)));
    ACCUM_FLT4 t01 = TO_ACCUM_TYPE(read_imagef(at_tex2d, smp_none, (int2)(DST_Y * 2 + 1, 0)));
    at_ar[0] = t00.x;
    at_ar[1] = t00.y;
    at_ar[2] = t00.z;
    at_ar[3] = t00.w;
    at_ar[4] = t01.x;
    at_ar[5] = t01.y;
    {
        ACCUM_FLT at = at_ar[0];
        ACCUM_FLT4 src0 = read_imagef(src_tensor_image2d, smp_none, (int2)((tile_id), (0) * shared_int4_1.x + (DST_Z)));
        I0 = at * src0;
        ACCUM_FLT4 src1 = read_imagef(src_tensor_image2d, smp_none, (int2)((tile_id), (1) * shared_int4_1.x + (DST_Z)));
        I1 = at * src1;
        ACCUM_FLT4 src2 = read_imagef(src_tensor_image2d, smp_none, (int2)((tile_id), (2) * shared_int4_1.x + (DST_Z)));
        I2 = at * src2;
        ACCUM_FLT4 src3 = read_imagef(src_tensor_image2d, smp_none, (int2)((tile_id), (3) * shared_int4_1.x + (DST_Z)));
        I3 = at * src3;
        ACCUM_FLT4 src4 = read_imagef(src_tensor_image2d, smp_none, (int2)((tile_id), (4) * shared_int4_1.x + (DST_Z)));
        I4 = at * src4;
        ACCUM_FLT4 src5 = read_imagef(src_tensor_image2d, smp_none, (int2)((tile_id), (5) * shared_int4_1.x + (DST_Z)));
        I5 = at * src5;
    }
    {
        ACCUM_FLT at = at_ar[1];
        ACCUM_FLT4 src0 = read_imagef(src_tensor_image2d, smp_none, (int2)((tile_id), (6) * shared_int4_1.x + (DST_Z)));
        I0 += at * src0;
        ACCUM_FLT4 src1 = read_imagef(src_tensor_image2d, smp_none, (int2)((tile_id), (7) * shared_int4_1.x + (DST_Z)));
        I1 += at * src1;
        ACCUM_FLT4 src2 = read_imagef(src_tensor_image2d, smp_none, (int2)((tile_id), (8) * shared_int4_1.x + (DST_Z)));
        I2 += at * src2;
        ACCUM_FLT4 src3 = read_imagef(src_tensor_image2d, smp_none, (int2)((tile_id), (9) * shared_int4_1.x + (DST_Z)));
        I3 += at * src3;
        ACCUM_FLT4 src4 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (10) * shared_int4_1.x + (DST_Z)));
        I4 += at * src4;
        ACCUM_FLT4 src5 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (11) * shared_int4_1.x + (DST_Z)));
        I5 += at * src5;
    }
    {
        ACCUM_FLT at = at_ar[2];
        ACCUM_FLT4 src0 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (12) * shared_int4_1.x + (DST_Z)));
        I0 += at * src0;
        ACCUM_FLT4 src1 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (13) * shared_int4_1.x + (DST_Z)));
        I1 += at * src1;
        ACCUM_FLT4 src2 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (14) * shared_int4_1.x + (DST_Z)));
        I2 += at * src2;
        ACCUM_FLT4 src3 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (15) * shared_int4_1.x + (DST_Z)));
        I3 += at * src3;
        ACCUM_FLT4 src4 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (16) * shared_int4_1.x + (DST_Z)));
        I4 += at * src4;
        ACCUM_FLT4 src5 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (17) * shared_int4_1.x + (DST_Z)));
        I5 += at * src5;
    }
    {
        ACCUM_FLT at = at_ar[3];
        ACCUM_FLT4 src0 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (18) * shared_int4_1.x + (DST_Z)));
        I0 += at * src0;
        ACCUM_FLT4 src1 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (19) * shared_int4_1.x + (DST_Z)));
        I1 += at * src1;
        ACCUM_FLT4 src2 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (20) * shared_int4_1.x + (DST_Z)));
        I2 += at * src2;
        ACCUM_FLT4 src3 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (21) * shared_int4_1.x + (DST_Z)));
        I3 += at * src3;
        ACCUM_FLT4 src4 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (22) * shared_int4_1.x + (DST_Z)));
        I4 += at * src4;
        ACCUM_FLT4 src5 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (23) * shared_int4_1.x + (DST_Z)));
        I5 += at * src5;
    }
    {
        ACCUM_FLT at = at_ar[4];
        ACCUM_FLT4 src0 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (24) * shared_int4_1.x + (DST_Z)));
        I0 += at * src0;
        ACCUM_FLT4 src1 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (25) * shared_int4_1.x + (DST_Z)));
        I1 += at * src1;
        ACCUM_FLT4 src2 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (26) * shared_int4_1.x + (DST_Z)));
        I2 += at * src2;
        ACCUM_FLT4 src3 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (27) * shared_int4_1.x + (DST_Z)));
        I3 += at * src3;
        ACCUM_FLT4 src4 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (28) * shared_int4_1.x + (DST_Z)));
        I4 += at * src4;
        ACCUM_FLT4 src5 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (29) * shared_int4_1.x + (DST_Z)));
        I5 += at * src5;
    }
    {
        ACCUM_FLT at = at_ar[5];
        ACCUM_FLT4 src0 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (30) * shared_int4_1.x + (DST_Z)));
        I0 += at * src0;
        ACCUM_FLT4 src1 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (31) * shared_int4_1.x + (DST_Z)));
        I1 += at * src1;
        ACCUM_FLT4 src2 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (32) * shared_int4_1.x + (DST_Z)));
        I2 += at * src2;
        ACCUM_FLT4 src3 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (33) * shared_int4_1.x + (DST_Z)));
        I3 += at * src3;
        ACCUM_FLT4 src4 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (34) * shared_int4_1.x + (DST_Z)));
        I4 += at * src4;
        ACCUM_FLT4 src5 = read_imagef(src_tensor_image2d, smp_none,
                                      (int2)((tile_id), (35) * shared_int4_1.x + (DST_Z)));
        I5 += at * src5;
    }
    ACCUM_FLT4 t0 = I1 + I2;
    ACCUM_FLT4 t1 = I3 + I4;
    FLT4 bias_val = read_imagef(biases_tex2d, smp_none, (int2)(DST_Z, 0));
    {
        FLT4 r0 = TO_FLT4(I0 + t0 + t1) + bias_val;
        {
            r0 = max(r0, (FLT) (0.0f));
        }
        write_imagef(dst_tensor_image2d, (int2)((tile_x), (tile_y) * shared_int4_0.w + (DST_Z)), r0);;
        tile_x++;
    }
    ACCUM_FLT4 t2 = I1 - I2;
    ACCUM_FLT4 t3 = I3 - I4;
    if (tile_x < shared_int4_0.y)
    {
        FLT4 r0 = TO_FLT4(t2 * At[7] + t3 * At[9]) + bias_val;
        {
            r0 = max(r0, (FLT) (0.0f));
        }
        write_imagef(dst_tensor_image2d, (int2)((tile_x), (tile_y) * shared_int4_0.w + (DST_Z)), r0);;
        tile_x++;
    }
    if (tile_x < shared_int4_0.y)
    {
        FLT4 r0 = TO_FLT4(t0 * At[13] + t1 * At[15]) + bias_val;
        {
            r0 = max(r0, (FLT) (0.0f));
        }
        write_imagef(dst_tensor_image2d, (int2)((tile_x), (tile_y) * shared_int4_0.w + (DST_Z)), r0);;
        tile_x++;
    }
    if (tile_x < shared_int4_0.y)
    {
        FLT4 r0 = TO_FLT4(t2 * At[19] + t3 * At[21] + I5) + bias_val;
        {
            r0 = max(r0, (FLT) (0.0f));
        }
        write_imagef(dst_tensor_image2d, (int2)((tile_x), (tile_y) * shared_int4_0.w + (DST_Z)), r0);;
        tile_x++;
    }
}